#include "update_kernels.h"
#include <device_launch_parameters.h>

#include <utility\cuda_utils.h>

namespace gc {
	__global__ void refitInteriors(int leafSize, LBvhLeaves leaves, LBvhInteriors interiors, int *visitCount) {
		
		int idx = blockIdx.x*blockDim.x + threadIdx.x;
		if (idx >= leafSize) return;

		int parent = leaves.parent[idx];
		//atomicMerge(interiors.bv, parent, leaves.bv, idx);
		//interiors.bv.merge(parent, leaves.bv.getBV(idx));
		if (interiors.lc[parent] == idx) {
			interiors.bv.setBV(parent, leaves.bv.getBV(idx));
		}
		else {
			interiors.bvR.setBV(parent, leaves.bv.getBV(idx));
		}

		// choose parent in bottom-up fashion. O(n)
		int current, left, right;
		while (atomicCAS(visitCount+parent,0,1)) {
			current = parent;
			interiors.bv.merge(current, interiors.bvR.getBV(current));
			if (current == 0) break; //root
			parent = interiors.parent[current];
			//atomicMerge(interiors.bv, parent, interiors.bv, current);
			//interiors.bv.merge(parent, interiors.bv.getBV(current));
			if (interiorIdx(interiors.lc[parent]) == current) {
				interiors.bv.setBV(parent, interiors.bv.getBV(current));
			}
			else {
				interiors.bvR.setBV(parent, interiors.bv.getBV(current));
			}
		}
		
	}
}